#include <device_atomic_functions.h>
#include "cuda_utils.h"

namespace gc {
	__device__ bool atomicMinf(float* address, float val) {
		int* address_as_i = (int*)address;
		int old = *address_as_i, assumed;
		if (*address <= val) return false;
		do {
			assumed = old;
			old = ::atomicCAS(address_as_i, assumed,
				__float_as_int(::fminf(val, __int_as_float(assumed))));
		} while (assumed != old);
		//return __int_as_float(old);
		return true;
	}

	__device__ bool atomicMaxf(float* address, float val) {
		int* address_as_i = (int*)address;
		int old = *address_as_i, assumed;
		if (*address >= val) return false;
		do {
			assumed = old;
			old = ::atomicCAS(address_as_i, assumed,
				__float_as_int(::fmaxf(val, __int_as_float(assumed))));
		} while (assumed != old);
		//return __int_as_float(old);
		return true;
	}

	__device__ bool atomicMinD(double* address, double val) {
		unsigned long long int* address_as_ull = (unsigned long long int*)address;
		unsigned long long int old = *address_as_ull, assumed;
		if (*address <= val) return false;
		do {
			assumed = old;
			old = ::atomicCAS(address_as_ull, assumed,
				__double_as_longlong(::fmin(val, __longlong_as_double(assumed))));
		} while (assumed != old);
		return true;
	}

	__device__ bool atomicMaxD(double* address, double val) {
		unsigned long long int* address_as_ull = (unsigned long long int*)address;
		unsigned long long int old = *address_as_ull, assumed;
		if (*address >= val) return false;
		do {
			assumed = old;
			old = ::atomicCAS(address_as_ull, assumed,
				__double_as_longlong(::fmax(val, __longlong_as_double(assumed))));
		} while (assumed != old);
		return true;
	}

	__host__ __device__ uint leftShift3f_32b(uint x) {
		if (x == (1 << 10)) --x;
		x = (x | (x << 16)) & 0b00000011000000000000000011111111;
		x = (x | (x << 8)) & 0b00000011000000001111000000001111;
		x = (x | (x << 4)) & 0b00000011000011000011000011000011;
		x = (x | (x << 2)) & 0b00001001001001001001001001001001;
		return x;
	}

	__host__ __device__ uint64 leftShift3f_64b(uint v) {
		uint64 x = v & 0x1fffff;
		x = (x | x << 32) & 0x1f00000000ffff;
		x = (x | x << 16) & 0x1f0000ff0000ff;
		x = (x | x << 8) & 0x100f00f00f00f00f;
		x = (x | x << 4) & 0x10c30c30c30c30c3;
		x = (x | x << 2) & 0x1249249249249249;
		return x;
	}

	__host__ __device__ uint morton3f_32b(float x, float y, float z) {
		uint ux = fmin(fmax(x * 1024.0f, 0.0f), 1023.0f);
		uint uy = fmin(fmax(y * 1024.0f, 0.0f), 1023.0f);
		uint uz = fmin(fmax(z * 1024.0f, 0.0f), 1023.0f);
		return (leftShift3f_32b(uz) << 2) |
			(leftShift3f_32b(uy) << 1) |
			leftShift3f_32b(ux);
	}

	__host__ __device__ uint64 morton3f_64b(float x, float y, float z) {
		uint ux = fmin(fmax(x * 2097152.0f, 0.0f), 2097151.0f);
		uint uy = fmin(fmax(y * 2097152.0f, 0.0f), 2097151.0f);
		uint uz = fmin(fmax(z * 2097152.0f, 0.0f), 2097151.0f);
		return (leftShift3f_64b(z) << 2) |
			(leftShift3f_64b(y) << 1) |
			leftShift3f_64b(x);
	}

}